__global__ void kernel_add(float *lhs, float *rhs, float *out, size_t elems_per_thread) {
    tops_dte_ctx_t ctx;
    //! use tops::dte_scope to initialize dte context
    tops::dte_scope s(ctx);
    auto thread_idx = threadIdx.z * (blockDim.x * blockDim.y) + threadIdx.y * blockDim.x + threadIdx.x;
    //! each thread does 1M times add
    static const size_t L1_BUFFER_LEN = 128;
    auto pos_start = thread_idx * elems_per_thread;
    auto pos_end = pos_start + elems_per_thread;
    //! L1 Buffer for lhs,rhs,out
    __valigned__ float l1_buf_lhs[L1_BUFFER_LEN];
    __valigned__ float l1_buf_rhs[L1_BUFFER_LEN];
    __valigned__ float l1_buf_out[L1_BUFFER_LEN];
    //! bind L1 Buffer to mdspan
    tops::mdspan l1_lhs(tops::Private, &l1_buf_lhs, L1_BUFFER_LEN);
    tops::mdspan l1_rhs(tops::Private, &l1_buf_rhs, L1_BUFFER_LEN);
    tops::mdspan l1_out(tops::Private, &l1_buf_out, L1_BUFFER_LEN);
    for (size_t i = pos_start; i < pos_end; i += L1_BUFFER_LEN) {
        //! bind L3 Buffer to mdspan
        tops::mdspan l3_lhs(tops::Global, lhs + i, L1_BUFFER_LEN);
        tops::mdspan l3_rhs(tops::Global, rhs + i, L1_BUFFER_LEN);
        tops::mdspan l3_out(tops::Global, out + i, L1_BUFFER_LEN);
        //! memcpy from L3 to L1.
        tops::memcpy(ctx, l1_lhs, l3_lhs);
        tops::memcpy(ctx, l1_rhs, l3_rhs);
        for (size_t j = 0; j < L1_BUFFER_LEN; j += tops::vlength<vfloat>()) {
            //! load data from L1 to vec register
            const auto &v_lhs = tops::vload<vfloat>(l1_buf_lhs + j);
            const auto &v_rhs = tops::vload<vfloat>(l1_buf_rhs + j);
            //! do add
            const auto &v_res = tops::vadd<vfloat>(v_lhs, v_rhs);
            //! store data to L1
            tops::vstore(v_res, l1_buf_out + j);
        }
        //! copy from L1 to L3
        tops::memcpy(ctx, l3_out, l1_out);
    }
}

void GCU_ADD(float * __restrict dev_lhs, float * __restrict dev_rhs, float * __restrict dev_out,  const int nr_elems) {
    static const size_t blocks = 1;
    static const size_t threads = 4;
    //! GCU ADD
    kernel_add<<<blocks, threads>>>(dev_lhs, dev_rhs, dev_out, nr_elems / (blocks * threads));
}
